{
 "cells": [
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "<h1><div align=\"center\">CUDA C/C++ で高速化されたアプリケーションの非同期ストリーミングと視覚的プロファイリング</div></h1>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "![CUDA](./images/CUDA_Logo.jpg)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "CUDA ツールキットは、GPU を活用する CUDA アプリケーションの開発をサポートする **NVIDIA Visual Profiler (nvvp)** という強力な GUI アプリケーションを標準搭載しています。アクセラレーテッド アプリケーションの視覚的なタイムラインを生成する **nvvp** では、CUDA API の呼び出し、カーネルの実行、メモリのアクティビティ、**CUDA ストリーム**の使用状況などの詳細な情報を参照できます。\n",
    "\n",
    "また、アクセラレーテッド アプリケーションの最適化に役立つインテリジェントな推奨案を確認できる各種分析ツールも揃っており、CUDA 開発者にとって、nvvp の習得は不可欠です。\n",
    "\n",
    "このラボでは、アクセラレーテッド アプリケーションの最適化に nvvp タイムラインを活用する方法を学習します。また、**アンマネージド メモリの割り当ておよび移行**、ホスト メモリの**ピン留め (ページロック)**、**非デフォルト CUDA ストリーム**といった、開発に役立つ CUDA の中級者向けプログラミング技術も学習します。\n",
    "\n",
    "ラボの最後に、簡単な N 体シミュレーターをアクセラレートおよび最適化するテストを行います。これは、コースで習得したスキルを実践できる内容になっています。正確さを維持しながらシミュレーターをアクセラレートできた方には、スキルを証明する認定が与えられます。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 前提条件\n",
    "\n",
    "このラボを効果的に活用するには、次のスキルを習得している必要があります。\n",
    "\n",
    "- CPU 関数の呼び出しと GPU カーネルの起動の両方を行う C/C++ プログラムを作成、コンパイル、実行する。\n",
    "- 実行構成を使用して、並列スレッド階層を制御する。\n",
    "- シリアル ループをリファクタリングして、ループの反復を GPU 上で並列実行する。\n",
    "- CUDA ユニファイド メモリの割り当てと解放を行う。\n",
    "- ページ フォールトとデータ移行に関してユニファイド メモリの動作を理解する。\n",
    "- 非同期的メモリ プリフェッチによって、ページ フォールトとデータ移行を削減する。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 目標\n",
    "\n",
    "このラボを完了すると、次のことができるようになります。\n",
    "\n",
    "- **NVIDIA Visual Profiler** (**nvvp**) を使用して、CUDA アプリケーションを視覚的にプロファイリングする。\n",
    "- nvvp を使用して、CUDA アプリケーションの最適化方法を特定し、活用する。\n",
    "- Utilize CUDA streams for concurrent kernel execution in accelerated applications.\n",
    "- (オプションの高度なコンテンツ) ピン留めされたメモリの割り当てなど、手動でメモリを割り当て、並行 CUDA ストリームにより非同期でデータを転送する。(**Optional Advanced Content**)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## NVIDIA Visual Profiler の設定\n",
    "\n",
    "[このリンク](/novnc)をクリックして nvvp に移動し、別のタブで nvvp セッションを開きます。パスワード `cuda` を使用して nvvp にアクセスします。次のセクションでは CUDA コードのプロファイルを作成します。\n",
    "\n",
    "**注: Windows のタッチ操作対応スクリーンのノート PC をご利用の方は、nvvp の使用時に問題が発生する場合があります。その場合は、Firefox Web ブラウザーを使用すると回避できます。** \n",
    "\n",
    "ワークスペースを使用するよう求められた場合は、選択された既定のワークスペースを承認するだけで、すぐに nvvp が自動で起動します。\n",
    "\n",
    "このラボ中に nvvp との接続がタイムアウトした場合、接続ボタンをクリックすれば再接続できます。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## nvvp でコードのリファクタリングを反復的に比較\n",
    "\n",
    "以下の一連の演習では、nvvp タイムラインの活用方法を学びます。既に習得している手法を使用して反復的に改善された各種プログラムのプロファイルを作成します。プロファイルを作成するごとに、次の反復時に役立つ情報をタイムラインで確認できます。これによりさまざまな CUDA プログラミング手法がアプリケーションのパフォーマンスにどのように影響するかを理解することができます。\n",
    "\n",
    "この一連の演習を完了すると、同時 CUDA ストリームを使用した新しい CUDA プログラミング手法を学んだり、アンマネージド CUDA メモリの割り当ておよびコピーの手法を使用したりする際に、nvvp タイムラインを活用できるようになります。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: コンパイル済み CUDA コードのタイムラインを分析する\n",
    "\n",
    "[`01-vector-add.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/01-vector-add/01-vector-add.cu) (<- リンクをクリックすると、ブラウザー内でソース ファイルを編集できます) には、アクセラレートされたベクトル加法アプリケーションが含まれます。すぐ下にあるコード実行セルを使用して (このコードを含むラボ内のすべてのコード実行セルは、`CTRL` を押しながらクリックしてください)、コンパイルおよび実行できます。成功したらメッセージが表示されます。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 2,
   "metadata": {},
   "outputs": [
    {
     "name": "stdout",
     "output_type": "stream",
     "text": [
      "Success! All values calculated correctly.\r\n"
     ]
    }
   ],
   "source": [
    "!nvcc -o vector-add-no-prefetch 01-vector-add/01-vector-add.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "アプリケーションを正常にコンパイルできたら、生成された実行可能ファイルを [nvvp](/novnc) で開いてタイムラインのウィンドウを最大化し、次の手順を実行します。\n",
    "\n",
    "- `addVectorsInto` カーネルの実行時間を示す、タイムラインのルーラーを作成します。\n",
    "- アプリケーションのタイムラインで、CPU ページ フォールトが発生している位置を特定します。[アプリケーションのソース コード](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/01-vector-add/01-vector-add.cu)内で、CPU ページ フォールトの発生原因となる場所を特定します。\n",
    "- タイムラインで、**[Data Migration (DtoH)]** (デバイスからホスト) のイベントを確認します。このイベントは、カーネル実行後に CPU ページ フォールトとほぼ同じタイミングで発生します。なぜこのイベントが、カーネル実行前の CPU ページ フォールト中ではなく、このタイミングで発生するのかを考えてみてください。\n",
    "- タイムラインから、GPU ページ フォールト、HtoD のデータ移行イベント、`addVectorsInto` カーネル実行の関係性を考えてみてください。ソース コードを参照し、これらのイベントがこのような形で発生する理由を説明できるでしょうか。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: 非同期プリフェッチを追加する\n",
    "\n",
    "[`01-vector-add-prefetch-solution.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/01-vector-add/solutions/01-vector-add-prefetch-solution.cu) は、前述のベクトル加法アプリケーションのリファクタリングを行い、`addVectorsInto` カーネルで必要な 3 つのベクトルが、カーネルの起動前にアクティブな GPU デバイスに非同期的にプリフェッチされるようにします ([`cudaMemPrefetchAsync`](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1ge8dc9199943d421bc8bc7f473df12e42) を使用)。ソース コードを開き、アプリケーション内でこれらの変更が適用された場所を特定します。\n",
    "\n",
    "変更を確認したら、すぐ下にあるコード実行セルを使用して、リファクタリングされたアプリケーションをコンパイルして実行します。成功したらメッセージが表示されます。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: プリフェッチした場合とプリフェッチしていない場合のタイムラインを比較する\n",
    "\n",
    "[`01-vector-add-prefetch-solution.cu`](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/01-vector-add/solutions/01-vector-add-prefetch-solution.cu) は、上述のベクトル加法アプリケーションのリファクタリングを行い、`addVectorsInto` カーネルで必要な 3 つのベクトルが、カーネルの起動前にアクティブな GPU デバイスに非同期的にプリフェッチされるようにします ([`cudaMemPrefetchAsync`](http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1ge8dc9199943d421bc8bc7f473df12e42) を使用)。ソース コードを開き、アプリケーション内でこれらの変更が適用された場所を特定します。\n",
    "\n",
    "変更を確認したら、すぐ下にあるコード実行セルを使用して、リファクタリングされたアプリケーションをコンパイルして実行します。成功したらメッセージが表示されます。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch sm_70 -o vector-add-prefetch 01-vector-add/solutions/01-vector-add-prefetch-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "[nvvp](/novnc) でコンパイルされた実行可能ファイルを開きます。ベクトル加算アプリケーションにプリフェッチを実装する前のセッションが開いたままになります。タイムライン ウィンドウを最大化し、次の手順を実行します。\n",
    "\n",
    "- `addVectorsInto` カーネルの実行時間を示す、タイムラインのルーラーを作成します。非同期的プリフェッチを追加する前の `addVectorsInto` カーネルと、実行時間を比較します。\n",
    "- タイムラインの *Runtime API* セクションで `cudaMemPrefetchAsync` を探します。\n",
    "- リファクタリング前のアプリケーションのタイムラインを参照すると、*Unified Memory* セクションで、ユニファイド メモリが必要になったため、カーネル実行中に *GPU page faults* のグループがいくつか発生したことを確認できます。プリフェッチ実装後に、これらのページ フォールトが発生しているか確認してください。\n",
    "- GPU ページ フォールトが解消された場合でも、引き続きデータをホストからデバイスに移行する必要があります。タイムラインの *Data Migration (HtoD)* セクションで、2 つのアプリケーション間でこれらの移行を比較します。`addVectorsInto` カーネルの実行に関連するデータ移行の回数、その実行に要した時間、実行されたタイミングを比較します。\n",
    "- 2 つのアプリケーションの全体的なランタイムを確認して比較します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "\n",
    "### 演習: カーネルの起動時に初期化するようリファクタしてプロファイリング\n",
    "\n",
    "ベクトル加法アプリケーションの前回のイテレーションでは、ベクトル データが CPU で初期化されるため、`addVectorsInto` カーネルで処理する前に GPU に移行する必要がありました。\n",
    "\n",
    "アプリケーション [01-init-kernel-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/02-init-kernel/solutions/01-init-kernel-solution.cu) の次のイテレーションでは、GPU でデータが並行して初期化されるようにアプリケーションがリファクタリングされています。\n",
    "\n",
    "GPU で初期化されるため、プリフェッチはベクトル加法処理の前ではなく、初期化の前に実行されます。ソース コードを確認して、これらの変更が適用された場所を特定します。\n",
    "\n",
    "変更を確認したら、すぐ下にあるコード実行セルを使用して、リファクタリングされたアプリケーションをコンパイルして実行します。成功したらメッセージが表示されます。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o init-kernel 02-init-kernel/solutions/01-init-kernel-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "コンパイルされた実行可能ファイルを nvvp の別のセッションで開き、次の手順を実行します。\n",
    "\n",
    "- タイムラインのルーラーを作成し、アプリケーションの全体的なランタイムと `addVectorsInto` カーネルのランタイムのほか、初期化カーネルのランタイムを測定します。アプリケーションと `addVectorsInto` のランタイムを前のバージョンのアプリケーションと比較し、どのように変更されたかを確認します。\n",
    "- タイムラインの *Compute* セクションを確認します。2 つのカーネル (`addVectorsInto` と初期化カーネル) のうち、GPU の処理時間の大部分を占有しているのはどちらかを特定してください。\n",
    "- 以下のうち、アプリケーションにどれが含まれているかを確認します。\n",
    "  - CPU Page Faults\n",
    "  - GPU Page Faults\n",
    "  - Data Migration (HtoD)\n",
    "  - Data Migration (DtoH)"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: データを非同期的にホストにプリフェッチするようリファクタしてプロファイリング\n",
    "\n",
    "現在、ベクトル加法アプリケーションは、ホスト上でベクトル加法カーネルの処理を確認しています。アプリケーション [01-prefetch-check-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/04-prefetch-check/solutions/01-prefetch-check-solution.cu) の次のリファクタリングでは、確認のためにデータを非同期的にホストにプリフェッチし直します。\n",
    "\n",
    "変更を確認したら、すぐ下にあるコード実行セルを使用して、リファクタリングされたアプリケーションをコンパイルして実行します。成功したらメッセージが表示されます。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o prefetch-to-host 04-prefetch-check/solutions/01-prefetch-check-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "nvvp で、新たにコンパイルされたこの実行可能ファイルを開いてタイムラインを最大化し、次の手順を実行します。\n",
    "\n",
    "- タイムラインの *Unified Memory* セクションで、CPU にプリフェッチを追加する前と後の *Data Migration (DtoH)* のイベントを比較します。\n",
    "  - CPU ページ フォールトの数を比較してください。\n",
    "  - DtoH の移行に要した合計時間を比較してください。\n",
    "- 次のセクションにスムーズに進むために、タイムラインの *Streams* セクションをチェックし、すべてのカーネル実行が *Default* ストリームで行われていることと、順番に処理されていることを確認します。次のセクションでは、ストリームについて学習します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 並行 CUDA ストリーム\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化した資料です。内容を確認してから、次のセクションのトピック詳細に進んでください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 1,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1LrLakdHiq5wtpI9SNqu2qPqVxDIeyjDizuSCmgwLOYQ/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1LrLakdHiq5wtpI9SNqu2qPqVxDIeyjDizuSCmgwLOYQ/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "CUDA プログラミングにおいて、**ストリーム**とは順番に実行される一連のコマンドのことを指します。CUDA アプリケーションでは、カーネルの実行と一部のメモリ転送は CUDA ストリームで行われます。ここまで CUDA ストリームには明示的に触れていませんでしたが、前の演習の nvvp タイムラインで確認したとおり、CUDA コードでは、*デフォルト ストリーム*内でカーネルを実行しています。\n",
    "\n",
    "CUDA プログラマは、デフォルト ストリームに加えて非デフォルト CUDA ストリームを作成して利用できます。これにより、別のストリームで複数のカーネルを同時に実行するなど、さまざまな処理を実行できます。複数のストリームを利用することで、アクセラレーテッド アプリケーションでの並列処理数が増え、アプリケーションの最適化の可能性がさらに広がります。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### CUDA ストリームの動作に関するルール\n",
    "\n",
    "CUDA ストリームの動作に関しては、効果的に利用するために覚えておきたい以下のようなルールがあります。\n",
    "\n",
    "- ストリーム内の処理は順番に実行される。\n",
    "- 非デフォルト ストリーム同士は、特定の順番で処理されるとは限らない。\n",
    "- デフォルト ストリームはブロッキングを行うので、他のすべてのストリームが完了するまで待ってから実行するだけでなく、その処理が完了するまで他のストリームが実行されないようにブロックする。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 非デフォルト CUDA ストリームの作成、利用、破棄\n",
    "\n",
    "次のコードは、非デフォルト CUDA ストリームを作成、利用、破棄する方法を示しています。非デフォルト CUDA ストリームで CUDA カーネルを起動するには、そのストリームを、実行構成のオプションである第 4 引数として渡す必要があります。これまでは、実行構成の最初の 2 つの引数のみを利用してきました。\n",
    "\n",
    "```cpp\n",
    "cudaStream_t stream;       // CUDA streams are of type `cudaStream_t`.\n",
    "cudaStreamCreate(&stream); // Note that a pointer must be passed to `cudaCreateStream`.\n",
    "\n",
    "someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>(); // `stream` is passed as 4th EC argument.\n",
    "\n",
    "cudaStreamDestroy(stream); // Note that a value, not a pointer, is passed to `cudaDestroyStream`.\n",
    "```\n",
    "\n",
    "このラボでは使用しませんが、実行構成の第 3 引数オプションも覚えておきましょう。これは、このカーネル起動のブロックごとに動的に割り当てられる、**共有メモリ** (現在は学習内容に含まれていない高度なトピック) のバイト数を指定する引数です。ブロックごとに共有メモリに割り当てられる既定のバイト数は 0 です。このラボでは、第 4 引数を示すために、第 3 引数の値として 0 を渡します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: デフォルト ストリームの動作を予測する\n",
    "\n",
    "The [01-print-numbers](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/05-stream-intro/01-print-numbers.cu) アプリケーションには、整数を受け入れて出力する、非常にシンプルな `printNumber` カーネルが含まれています。このカーネルは、単一ブロック内の 1 つのスレッドでのみ処理されますが、for ループで 5 回実行され、起動ごとに for ループの反復回数を渡します。\n",
    "\n",
    "以下のコード実行ブロックを使用して、[01-print-numbers](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/05-stream-intro/01-print-numbers.cu) をコンパイルして実行します。`0` から `4` までの数値が出力されることがわかります。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o print-numbers 05-stream-intro/01-print-numbers.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "既定ではカーネルがデフォルト ストリームで実行されますが、`print-numbers` プログラムの 5 回の起動は順番に行われるか、並列で行われるかを予想してください。デフォルト ストリームの 2 つの特徴から、答えを導き出してください。nvvp の新しいセッションで実行可能ファイルを開いてタイムラインを最大化してから、カーネルの起動に注目して答えを確認してください。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: 並行 CUDA ストリームを実装する\n",
    "\n",
    "5 つのカーネル起動はすべて同一のストリーム内で実行されるため、カーネルは予想どおり順番に実行されます。また、デフォルト ストリームがブロッキングを行うため、1 つのカーネルの起動が完了してから次のカーネルが起動することもわかりました。\n",
    "\n",
    "[01-print-numbers](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/05-stream-intro/01-print-numbers.cu) をリファクタリングして、各カーネルがそれぞれの非デフォルト ストリームで起動するようにします。作成したストリームは、不要になったら必ず破棄してください。以下のコード実行セルを使用して、リファクタリングされたコードをコンパイルして実行します。引き続き `0` から `4` までの数値が出力されますが、必ずしも昇順とは限りません。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/05-stream-intro/solutions/01-print-numbers-solution.cu) を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o print-numbers-in-streams 05-stream-intro/01-print-numbers.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "5 つのカーネルを起動するのに、非デフォルトの 5 つのストリームを使用するように設定されました。今度は順番に実行されるでしょうか、それとも並列に実行されるでしょうか。ストリームの特性だけでなく、`printNumber` カーネルの単純さも考慮してください。つまり、並列で実行されると予測した場合に、カーネルが完全に同時に処理されるような実行速度であるかどうかを考えます。\n",
    "\n",
    "予測を立てたら、リファクタリングしたプログラムの実行可能ファイルを使用して新しい nvvp セッションを開き、タイムラインを最大化して、実際の動作を確認します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: 同時データ初期化カーネルのストリームを使用する\n",
    "\n",
    "使用しているベクトル加法アプリケーション [01-prefetch-check-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/04-prefetch-check/solutions/01-prefetch-check-solution.cu) は、初期化カーネルを 3 回起動します。`vectorAdd` カーネルで、初期化が必要な 3 つのベクトルで 1 回ずつ起動します。このアプリケーションをリファクタリングし、3 回の初期化カーネル起動が、それぞれの非デフォルト ストリームで行われるようにします。以下のコード実行セルでコンパイルして実行すると、成功メッセージが出力されます。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/06-stream-init/solutions/01-stream-init-solution.cu) を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o init-in-streams 04-prefetch-check/solutions/01-prefetch-check-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "コンパイルされたバイナリを nvvp で開いてタイムラインを最大化してから、3 回の初期化カーネル起動がそれぞれの非デフォルト ストリームで実行され、ほぼ同時に処理されていることを確認します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## まとめ\n",
    "\n",
    "ここまでで、次のことができるようになりました。\n",
    "\n",
    "- **NVIDIA Visual Profiler** (**nvvp**) を使用して、GPU 高速化 CUDA アプリケーションの視覚的なプロファイルを作成する。\n",
    "- nvvp を使用して、GPU 高速化 CUDA アプリケーションの最適化方法を特定し、活用する。\n",
    "- CUDA ストリームを利用し、CUDA アプリケーションで並行にカーネル実行を行う。\n",
    "\n",
    "CPU 専用アプリケーションをアクセラレートしてそのアクセラレーテッド アプリケーションを最適化するための基本的なツールや手法を習得しました。最後の演習では、これまでに学んだことをすべて応用して [N 体](https://en.wikipedia.org/wiki/N-body_problem)シミュレーターをアクセラレートします。このシミュレーターでは、重力に従って相互作用するオブジェクトのグループの個々の動きを予測します。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 最後の演習: N 体シミュレーターを高速化して最適化する\n",
    "\n",
    "[N 体](https://en.wikipedia.org/wiki/N-body_problem) シミュレーターは、重力に従って相互作用するオブジェクト グループの個々の動きを予測します。[01-nbody.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/09-nbody/01-nbody.cu) には、3 次元空間を移動する物体のために、シンプルながら実用的な N 体シミュレーターが用意されています。このアプリケーションにコマンド ライン引数を渡して、システム内の物体の数を指定できます。\n",
    "\n",
    "4,096 個の物体を対象とする CPU 専用の形式の場合、このアプリケーションは、システム内の物体間での相互作用を 1 秒につき約 3,000 万回計算できます。ここでは、次のことを行います。\n",
    "\n",
    "- シミュレーションの正確さを維持しながら、GPU でプログラムを高速化する\n",
    "- 反復的な処理を行ってシミュレーターを最適化し、4,096 個の物体を処理しながら、1 秒間に 300 億回の相互作用を計算できるようにする `(2<<11)`\n",
    "- 反復的な処理を行ってシミュレーターを最適化し、最大 65,000 個の物体を処理しながら、1 秒間に 3250 億回の相互作用を計算できるようにする `(2<<15)`\n",
    "\n",
    "**この演習を完了したら、このノートを開く前のブラウザーのページに戻り、[Assess] ボタンをクリックします。アプリケーション精度を維持したまま、上記の仕様に合わせてアクセラレートすることに成功した場合、「Fundamentals of Accelerated Computing with CUDA C/C++ (CUDA C/C++ によるアクセラレーテッド コンピューティングの基礎)」の認定が与えられます。**"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 課題に取り組む際のポイント\n",
    "\n",
    "課題を始める前に、以下の検討事項をお読みください。\n",
    "\n",
    "- 最初のリファクタリングの際には、アプリケーションのロジック (特に `bodyForce` 関数) をほとんど変えないようにします。できるだけ簡単に高速化することを心がけてください。\n",
    "- `randomizeBodies` 関数は、GPU デバイスに対応していない `rand` 関数を使用しているため、高速化できません。`randomizeBodies` はホスト関数ですので、手を加えないでください。\n",
    "- コードベースの `main` には `for` ループを含めます。これは、`bodyForce` で計算された物体間の力をシステム内の物体の位置に統合するためです。この統合は、`bodyForce` の実行後に開始し、次の `bodyForce` の呼び出しの前に完了させる必要があります。並列化の方法と位置を決定する際に、この点を考慮してください。\n",
    "- プロファイル駆動の反復手法を使用します。\n",
    "- コードにエラー処理を追加する必要はありませんが、コードを正しく動作させるために、追加しておくと役に立ちます。\n",
    "\n",
    "それでは、頑張ってください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o nbody 09-nbody/01-nbody.cu"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!./nbody 11 # This argument is passed as `N` in the formula `2<<N`, to determine the number of bodies in the system"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvprof ./nbody"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "## 追加コンテンツ\n",
    "\n",
    "以下のセクションは、余裕のある方を対象とした追加の課題です。ここでは、手動によるメモリ管理、非デフォルト ストリームを使用したカーネル実行とメモリ コピーの同時処理を含む、中級レベルの手法を紹介します。\n",
    "\n",
    "以下の手法を理解したら、それを使用して N 体シミュレーションをさらに最適化してみましょう。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## 手動でのメモリの割り当てとコピー\n",
    "\n",
    "`cudaMallocManaged` と `cudaMemPrefetchAsync` はパフォーマンスが高く、メモリ移行を大幅に簡略化できますが、手動でメモリを割り当てる方が適していることもあります。特に、データ アクセスがデバイスまたはホストからに限定されるなど、自動のオンデマンド移行が必要なくデータ移行のコストを再分配できる場合に当てはまります。\n",
    "\n",
    "さらに、手動のメモリ管理では、非デフォルト ストリームを使用して、データ転送と計算処理を同時に実行できます。このセクションでは、手動のメモリの割り当ておよびコピーの基本的な手法を学習した後、これらの手法を応用してデータのコピーと計算処理を同時に実行します。\n",
    "\n",
    "手動でのメモリ管理に使用する CUDA 関数を次に示します。\n",
    "\n",
    "- `cudaMalloc` は、アクティブな GPU にメモリを直接割り当てます。これにより、GPU ページ フォールトをすべて防止できます。その代わり、返されるポインターにホスト コードでアクセスすることはできません。\n",
    "- `cudaMallocHost` は CPU にメモリを直接割り当てます。さらに、メモリの「ピン留め」 (ページ ロック) も行います。これにより、GPU 内外へのメモリの非同期コピーが可能になります。ピン留めされたメモリが多すぎると CPU のパフォーマンスに悪影響が及ぶおそれがあるため、目的が決まっている場合のみ使用してください。ピン留めされたメモリは、`cudaFreeHost` を使用して解放する必要があります。\n",
    "- `cudaMemcpy` は、ホストからデバイスまたはデバイスからホストに、メモリを (転送ではなく) コピーします。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 手動でのメモリ管理の例\n",
    "\n",
    "上記の CUDA API 呼び出しの使用例を示すコードは次のとおりです。\n",
    "\n",
    "```cpp\n",
    "int *host_a, *device_a;        // Define host-specific and device-specific arrays.\n",
    "cudaMalloc(&device_a, size);   // `device_a` is immediately available on the GPU.\n",
    "cudaMallocHost(&host_a, size); // `host_a` is immediately available on CPU, and is page-locked, or pinned.\n",
    "\n",
    "initializeOnHost(host_a, N);   // No CPU page faulting since memory is already allocated on the host.\n",
    "\n",
    "// `cudaMemcpy` takes the destination, source, size, and a CUDA-provided variable for the direction of the copy.\n",
    "cudaMemcpy(device_a, host_a, size, cudaMemcpyHostToDevice);\n",
    "\n",
    "kernel<<<blocks, threads, 0, someStream>>>(device_a, N);\n",
    "\n",
    "// `cudaMemcpy` can also copy data from device to host.\n",
    "cudaMemcpy(host_a, device_a, size, cudaMemcpyDeviceToHost);\n",
    "\n",
    "verifyOnHost(host_a, N);\n",
    "\n",
    "cudaFree(device_a);\n",
    "cudaFreeHost(host_a);          // Free pinned memory like this.\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### 演習: ホストおよびデバイスのメモリを手動で割り当てる\n",
    "\n",
    "ベクトル加法アプリケーション [01-stream-init-solution](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/06-stream-init/solutions/01-stream-init-solution.cu) の直近の反復では、`cudaMallocManaged` を使用して、初期化カーネルがデバイスで使用するマネージド メモリ、ベクトル加法カーネルがデバイスで使用するメモリ、ホストが使用するメモリの順に割り当てます (メモリは検証用に自動転送されます)。これは堅実な手法ですが、手動によるメモリの割り当てとコピーを試し、アプリケーションのパフォーマンスへの影響を確認することも大事です。\n",
    "\n",
    "`cudaMallocManaged` を使用**しない**ように [01-stream-init-solution](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/06-stream-init/solutions/01-stream-init-solution.cu) アプリケーションをリファクタリングします。そのためには、次の手順を実行する必要があります。\n",
    "\n",
    "- `cudaMallocManaged` の呼び出しを `cudaMalloc` に置き換えます。\n",
    "- ホストでの検証に使用する追加のベクトルを作成します。`cudaMalloc` が割り当てるメモリをホストで利用できないため、これを行う必要があります。このホストのベクトル割り当てには、`cudaMallocHost` を使用します。\n",
    "- `addVectorsInto` カーネルを実行したら、`cudaMemcpy` を使用して、加算結果とベクトルを、`cudaMallocHost` で作成したホストのベクトルにコピーします。\n",
    "- `cudaFreeHost` を使用して、`cudaMallocHost` によって割り当てられたメモリを解放します。\n",
    "\n",
    "行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/07-manual-malloc/solutions/01-manual-malloc-solution.cu)を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o vector-add-manual-alloc 06-stream-init/solutions/01-stream-init-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "リファクタリングが完了したら、新しい nvvp セッションで実行可能ファイルを開き、タイムラインで次の点を確認します。\n",
    "\n",
    "- タイムラインの *Unified Memory* セクションが表示されていないことを確認します。\n",
    "- このタイムラインと、前のリファクタリング時のタイムラインを比較します。タイムラインのルーラーを使用して、現在のアプリケーションの `cudaMalloc` と前のアプリケーションの `cudaMallocManaged` のランタイムを比較してください。- 現在のアプリケーションの初期化カーネル処理の開始タイミングが、前回のイテレーションよりも遅いことを確認します。タイムライン上では、`cudaMallocHost` の所要時間の違いがわかります。これにより、メモリの転送とメモリのコピーの違いが明確になります。今回のようにメモリをコピーすると、データはシステム内の 2 か所に存在することになります。今回のケースでは、ホスト専用の 4 つ目のベクトルを割り当てた際に、前回のイテレーションで 3 つのベクトルのみを割り当てた場合と比べて、パフォーマンスにわずかに負担が生じています。"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "---\n",
    "## ストリームを使用したデータ移行とコード実行の同時処理\n",
    "\n",
    "次のスライドは、このセクションの概要を視覚化した資料です。内容を確認してから、次のセクションのトピック詳細に進んでください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": 6,
   "metadata": {},
   "outputs": [
    {
     "data": {
      "text/html": [
       "\n",
       "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1sqgvems_bMQrHGD3IMeRdIwwTn0nlWAAKbFaoTXL1Ag/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
      ],
      "text/plain": [
       "<IPython.core.display.HTML object>"
      ]
     },
     "metadata": {},
     "output_type": "display_data"
    }
   ],
   "source": [
    "%%HTML\n",
    "\n",
    "<div align=\"center\"><iframe src=\"https://docs.google.com/presentation/d/1sqgvems_bMQrHGD3IMeRdIwwTn0nlWAAKbFaoTXL1Ag/embed?start=false&loop=false&delayms=3000\" frameborder=\"0\" width=\"900\" height=\"550\" allowfullscreen=\"true\" mozallowfullscreen=\"true\" webkitallowfullscreen=\"true\"></iframe></div>"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "`cudaMemcpy` 以外に、`cudaMemcpyAsync` では、`cudaMallocHost` を使用して割り当てることでホストのメモリがピン留めされている限り、ホストからデバイスへも、デバイスからホストへも、メモリを非同期的にコピーできます。\n",
    "\n",
    "カーネル実行と同様に、`cudaMemcpyAsync` は、既定でホストに対してのみ非同期処理を行います。既定では、デフォルト ストリームで実行されるため、GPU で行われる他の CUDA 処理についてはブロッキング処理となります。ただし、`cudaMemcpyAsync` 関数は、第 5 引数のオプションで、非デフォルト ストリームを使用します。CPU 依存のストリームを渡すことで、メモリ転送を、他の非デフォルト ストリームでの CUDA 処理と並行して実行できます。\n",
    "\n",
    "ピン留めされたホスト メモリ、非デフォルト ストリームでの非同期的なメモリ コピー、非デフォルト ストリームでのカーネル実行を組み合わせて、メモリ転送とカーネル実行を同時に処理するのが、よく使われる便利なパターンです。\n",
    "\n",
    "次の例では、メモリ コピー全体の完了を待ってからカーネル処理を開始するのではなく、各非デフォルト ストリームで必要なデータのセグメントをコピーおよび処理します。この手法では、データの一部に対する処理と、後で処理されるセグメントのメモリ転送を同時に開始できます。その際、次に示すように、セグメントごとの処理回数の値と、配列内のオフセットの位置の算出に特に注意を払ってください。\n",
    "\n",
    "```cpp\n",
    "int N = 2<<24;\n",
    "int size = N * sizeof(int);\n",
    "\n",
    "int *host_array;\n",
    "int *device_array;\n",
    "\n",
    "cudaMallocHost(&host_array, size);               // Pinned host memory allocation.\n",
    "cudaMalloc(&device_array, size);                 // Allocation directly on the active GPU device.\n",
    "\n",
    "initializeData(host_array, N);                   // Assume this application needs to initialize on the host.\n",
    "\n",
    "const int numberOfSegments = 4;                  // This example demonstrates slicing the work into 4 segments.\n",
    "int segmentN = N / numberOfSegments;             // A value for a segment's worth of `N` is needed.\n",
    "size_t segmentSize = size / numberOfSegments;    // A value for a segment's worth of `size` is needed.\n",
    "\n",
    "// For each of the 4 segments...\n",
    "for (int i = 0; i < numberOfSegments; ++i)\n",
    "{\n",
    "  // Calculate the index where this particular segment should operate within the larger arrays.\n",
    "  segmentOffset = i * segmentN;\n",
    "\n",
    "  // Create a stream for this segment's worth of copy and work.\n",
    "  cudaStream_t stream;\n",
    "  cudaStreamCreate(&stream);\n",
    "  \n",
    "  // Asynchronously copy segment's worth of pinned host memory to device over non-default stream.\n",
    "  cudaMemcpyAsync(&device_array[segmentOffset],  // Take care to access correct location in array.\n",
    "                  &host_array[segmentOffset],    // Take care to access correct location in array.\n",
    "                  segmentSize,                   // Only copy a segment's worth of memory.\n",
    "                  cudaMemcpyHostToDevice,\n",
    "                  stream);                       // Provide optional argument for non-default stream.\n",
    "                  \n",
    "  // Execute segment's worth of work over same non-default stream as memory copy.\n",
    "  kernel<<<number_of_blocks, threads_per_block, 0, stream>>>(&device_array[segmentOffset], segmentN);\n",
    "  \n",
    "  // `cudaStreamDestroy` will return immediately (is non-blocking), but will not actually destroy stream until\n",
    "  // all stream operations are complete.\n",
    "  cudaStreamDestroy(stream);\n",
    "}\n",
    "```"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "### Exercise: 演習: カーネル実行とホストへのメモリ コピーを同時に処理する\n",
    "\n",
    "ベクトル加法アプリケーション [01-manual-malloc-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/07-manual-malloc/solutions/01-manual-malloc-solution.cu) の直近のイテレーションでは、GPU に対してすべてのベクトル加法処理を実行してから、検証用にメモリをホストにコピーしています。\n",
    "\n",
    "[01-manual-malloc-solution.cu](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/07-manual-malloc/solutions/01-manual-malloc-solution.cu) をリファクタリングして、非デフォルト ストリームで 4 つのセグメントに分けてベクトル加法を実行すると、すべてのベクトル加法処理が完了する前に、非同期でメモリ コピーを開始できます。行き詰まったときは、[解決策](../../../../../edit/tasks/task1/task/03_AC_STREAMS_NVVP/08-overlap-xfer/solutions/01-overlap-xfer-solution.cu) を参照してください。"
   ]
  },
  {
   "cell_type": "code",
   "execution_count": null,
   "metadata": {
    "collapsed": true
   },
   "outputs": [],
   "source": [
    "!nvcc -arch=sm_70 -o vector-add-manual-alloc 07-manual-malloc/solutions/01-manual-malloc-solution.cu -run"
   ]
  },
  {
   "cell_type": "markdown",
   "metadata": {},
   "source": [
    "リファクタリングが完了したら、新しい nvvp セッションで実行可能ファイルを開き、タイムラインで次の点を確認します。\n",
    "\n",
    "- デバイスからホストへのメモリ転送の開始時点で、カーネルの処理がすべて完了しているかどうかを確認します。\n",
    "- 4 つのメモリ コピー セグメント自体が同時実行されていないことを確認します。別々の非デフォルト ストリームにあっても、一度に同時実行できるのは、特定方向 (この場合は DtoH) のメモリ転送だけです。この方法では、早い段階で転送を開始できるため、パフォーマンスが向上し、当然ながら、単純な加算処理に比べて細かい処理が少なくなります。"
   ]
  }
 ],
 "metadata": {
  "kernelspec": {
   "display_name": "Python 2",
   "language": "python",
   "name": "python2"
  },
  "language_info": {
   "codemirror_mode": {
    "name": "ipython",
    "version": 2
   },
   "file_extension": ".py",
   "mimetype": "text/x-python",
   "name": "python",
   "nbconvert_exporter": "python",
   "pygments_lexer": "ipython2",
   "version": "2.7.12"
  }
 },
 "nbformat": 4,
 "nbformat_minor": 1
}
